{
 "cells": [
  {
   "attachments": {},
   "cell_type": "markdown",
   "id": "2e486d95-6c70-4c68-b7bb-83b97579e316",
   "metadata": {},
   "source": [
    "# Code Generation with Dynamic Shape\n",
    "**Large Language Models are dynamic**\n",
    "![image.png](./img/DynamicTuning.png)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "7e250c4a-d27f-4426-b67a-a29b1962cf94",
   "metadata": {},
   "outputs": [],
   "source": [
    "import bitblas\n",
    "import torch\n",
    "\n",
    "# enabling debug output\n",
    "\n",
    "bitblas.set_log_level(\"Debug\")"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 1,
   "id": "4258ce53-a8e7-4cd3-80bc-7e71518bb9e6",
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "2024-10-25 08:25:39 [BitBLAS:INFO]: Auto detected target: nvidia/geforce-rtx-4090\n",
      "\n",
      "template <typename T1, typename T2, bool isSigned = false>\n",
      "__device__ void decode_i4b_to_f16(T1 *_i4s, T2 *B_local_decode, const int N = 8)\n",
      "{\n",
      "    uint *h = reinterpret_cast<uint *>(B_local_decode);\n",
      "\n",
      "    static constexpr uint immLut = (0xf0 & 0xcc) | 0xaa;\n",
      "    static constexpr uint BOTTOM_MASK = 0x000f000f;\n",
      "    static constexpr uint FP16_TOP_MAGIC_NUM = 0x64006400;\n",
      "    static constexpr uint MEDIAN_NUM = isSigned ? 0x64086408 : 0x64006400;\n",
      "    uint const i4s = *reinterpret_cast<uint *>(_i4s);\n",
      "#pragma unroll\n",
      "    for (int i = 0; i < (N / 2); i++)\n",
      "    {\n",
      "\n",
      "        asm volatile(\"lop3.b32 %0, %1, %2, %3, %4;\\n\"\n",
      "                     : \"=r\"(h[i])\n",
      "                     : \"r\"(i4s >> (4 * i)), \"n\"(BOTTOM_MASK), \"n\"(FP16_TOP_MAGIC_NUM), \"n\"(immLut));\n",
      "        asm volatile(\"sub.f16x2 %0, %1, %2;\\n\" : \"=r\"(h[i]) : \"r\"(h[i]), \"r\"(MEDIAN_NUM));\n",
      "    }\n",
      "}\n",
      "\n",
      "template <typename T1, typename T2>\n",
      "__device__ void decode_i4s_to_f16(T1 *_i4s, T2 *B_local_decode, const int N = 8)\n",
      "{\n",
      "    decode_i4b_to_f16<T1, T2, true>(_i4s, B_local_decode, N);\n",
      "}\n",
      "\n",
      "template <typename T1, typename T2>\n",
      "__device__ void decode_i4u_to_f16(T1 *_i4u, T2 *B_local_decode, const int N = 8)\n",
      "{\n",
      "    decode_i4b_to_f16<T1, T2, false>(_i4u, B_local_decode, N);\n",
      "}\n",
      "#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 530)\n",
      "\n",
      "typedef unsigned short uint16_t;\n",
      "typedef unsigned char uint8_t;\n",
      "typedef signed char int8_t;\n",
      "typedef int int32_t;\n",
      "typedef unsigned long long uint64_t;\n",
      "typedef unsigned int uint32_t;\n",
      "\n",
      "#define TVM_FORCE_INLINE inline __attribute__((always_inline))\n",
      "#define TVM_XINLINE TVM_FORCE_INLINE __device__ __host__\n",
      "#define TVM_ALIGNED(x) __attribute__ ((aligned(x)))\n",
      "#define TVM_HALF_OPERATOR(RTYPE, OP)                              \\\n",
      "  TVM_XINLINE RTYPE operator OP (half a, half b) {                \\\n",
      "    return RTYPE(float(a) OP float(b));                           \\\n",
      "  }                                                               \\\n",
      "  template<typename T>                                            \\\n",
      "  TVM_XINLINE RTYPE operator OP (half a, T b) {                   \\\n",
      "    return RTYPE(float(a) OP float(b));                           \\\n",
      "  }                                                               \\\n",
      "  template<typename T>                                            \\\n",
      "  TVM_XINLINE RTYPE operator OP (T a, half b) {                   \\\n",
      "    return RTYPE(float(a) OP float(b));                           \\\n",
      "  }\n",
      "\n",
      "#define TVM_HALF_ASSIGNOP(AOP, OP)                                \\\n",
      "  template<typename T>                                            \\\n",
      "  TVM_XINLINE half operator AOP (const T& a) {                    \\\n",
      "    return *this = half(float(*this) OP float(a));                \\\n",
      "  }                                                               \\\n",
      "  template<typename T>                                            \\\n",
      "  TVM_XINLINE half operator AOP (const volatile T& a) volatile {  \\\n",
      "    return *this = half(float(*this) OP float(a));                \\\n",
      "  }\n",
      "\n",
      "class TVM_ALIGNED(2) half {\n",
      " public:\n",
      "  uint16_t half_;\n",
      "\n",
      "  static TVM_XINLINE half Binary(uint16_t value) {\n",
      "    half res;\n",
      "    res.half_ = value;\n",
      "    return res;\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE half() {}\n",
      "\n",
      "  TVM_XINLINE half(const float& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const double& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const int8_t& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const uint8_t& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const int32_t& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const uint32_t& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const long long& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const uint64_t& value) { constructor(value); }\n",
      "\n",
      "  TVM_XINLINE operator float() const {                          \\\n",
      "    return float(half2float(half_));                            \\\n",
      "  }                                                             \\\n",
      "  TVM_XINLINE operator float() const volatile {                 \\\n",
      "    return float(half2float(half_));                            \\\n",
      "  }\n",
      "\n",
      "\n",
      "  TVM_HALF_ASSIGNOP(+=, +)\n",
      "  TVM_HALF_ASSIGNOP(-=, -)\n",
      "  TVM_HALF_ASSIGNOP(*=, *)\n",
      "  TVM_HALF_ASSIGNOP(/=, /)\n",
      "\n",
      "  TVM_XINLINE half operator+() {\n",
      "    return *this;\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE half operator-() {\n",
      "    return half(-float(*this));\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE half operator=(const half& a) {\n",
      "    half_ = a.half_;\n",
      "    return a;\n",
      "  }\n",
      "\n",
      "  template<typename T>\n",
      "  TVM_XINLINE half operator=(const T& a) {\n",
      "    return *this = half(a);\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE half operator=(const half& a) volatile {\n",
      "    half_ = a.half_;\n",
      "    return a;\n",
      "  }\n",
      "\n",
      "  template<typename T>\n",
      "  TVM_XINLINE half operator=(const T& a) volatile {\n",
      "    return *this = half(a);\n",
      "  }\n",
      "\n",
      " private:\n",
      "  union Bits {\n",
      "    float f;\n",
      "    int32_t si;\n",
      "    uint32_t ui;\n",
      "  };\n",
      "\n",
      "  static int const fp16FractionBits = 10;\n",
      "  static int const fp32FractionBits = 23;\n",
      "  static int32_t const fp32FractionMask = ~(~0u << fp32FractionBits);   // == 0x7fffff\n",
      "  static int32_t const fp32HiddenBit = 1 << fp32FractionBits;   // == 0x800000\n",
      "  static int const shift = fp32FractionBits - fp16FractionBits;   // == 13\n",
      "  static int const shiftSign = 16;\n",
      "  static int32_t const expAdjust = 127 - 15;   // exp32-127 = exp16-15, so exp16 = exp32 - (127-15)\n",
      "\n",
      "  static int32_t const infN = 0x7F800000;   // flt32 infinity\n",
      "  static int32_t const maxN = 0x477FFFFF;   // max flt32 that's a flt16 normal after >> by shift\n",
      "  static int32_t const minN = 0x38800000;   // min flt16 normal as a flt32\n",
      "  static int32_t const maxZ = 0x33000000;   // max fp32 number that's still rounded to zero in fp16\n",
      "  static int32_t const signN = 0x80000000;  // flt32 sign bit\n",
      "\n",
      "  static int32_t const infC = infN >> shift;\n",
      "  static int32_t const nanN = (infC + 1) << shift;   // minimum flt16 nan as a flt32\n",
      "  static int32_t const maxC = maxN >> shift;\n",
      "  static int32_t const minC = minN >> shift;\n",
      "  static int32_t const signC = signN >> shiftSign;  // flt16 sign bit\n",
      "\n",
      "  static int32_t const mulN = 0x52000000;  // (1 << 23) / minN\n",
      "  static int32_t const mulC = 0x33800000;  // minN / (1 << (23 - shift))\n",
      "\n",
      "  static int32_t const subC = 0x003FF;  // max flt32 subnormal down shifted\n",
      "  static int32_t const norC = 0x00400;  // min flt32 normal down shifted\n",
      "\n",
      "  static int32_t const maxD = infC - maxC - 1;\n",
      "  static int32_t const minD = minC - subC - 1;\n",
      "\n",
      "  TVM_XINLINE uint16_t float2half(const float& value) const {\n",
      "    Bits v;\n",
      "    v.f = value;\n",
      "    uint32_t sign = v.si & signN;    // grab sign bit\n",
      "    v.si ^= sign;                    // clear sign bit from v\n",
      "    sign >>= shiftSign;              // logical shift sign to fp16 position\n",
      "\n",
      "    if (v.si <= maxZ) {\n",
      "      // Handle eventual zeros here to ensure\n",
      "      // vshift will not exceed 32 below.\n",
      "      v.ui = 0;\n",
      "    } else if (v.si < minN) {\n",
      "      // Handle denorms\n",
      "      uint32_t exp32 = v.ui >> fp32FractionBits;\n",
      "      int32_t exp16 = exp32 - expAdjust;\n",
      "      // If exp16 == 0 (just into the denorm range), then significant should be shifted right 1.\n",
      "      // Smaller (so negative) exp16 values should result in greater right shifts.\n",
      "      uint32_t vshift = 1 - exp16;\n",
      "      uint32_t significand = fp32HiddenBit | (v.ui & fp32FractionMask);\n",
      "      v.ui = significand >> vshift;\n",
      "      v.ui += (v.ui & 0x3fff) != 0x1000 || (significand & 0x7ff) ? 0x1000 : 0;\n",
      "    } else if (v.si <= maxN) {\n",
      "      // Handle norms\n",
      "      v.ui += (v.ui & 0x3fff) != 0x1000 ? 0x1000 : 0;\n",
      "      v.ui -= expAdjust << fp32FractionBits;\n",
      "    } else if (v.si <= infN) {\n",
      "      v.si = infN;\n",
      "    } else if (v.si < nanN) {\n",
      "      v.si = nanN;\n",
      "    }\n",
      "\n",
      "    v.ui >>= shift;\n",
      "    return sign | (v.ui & 0x7fff);\n",
      "  }\n",
      "\n",
      "  // Same as above routine, except for addition of volatile keyword\n",
      "  TVM_XINLINE uint16_t float2half(\n",
      "    const volatile float& value) const volatile {\n",
      "    Bits v;\n",
      "    v.f = value;\n",
      "    uint32_t sign = v.si & signN;    // grab sign bit\n",
      "    v.si ^= sign;                    // clear sign bit from v\n",
      "    sign >>= shiftSign;              // logical shift sign to fp16 position\n",
      "\n",
      "    if (v.si <= maxZ) {\n",
      "      // Handle eventual zeros here to ensure\n",
      "      // vshift will not exceed 32 below.\n",
      "      v.ui = 0;\n",
      "    } else if (v.si < minN) {\n",
      "      // Handle denorms\n",
      "      uint32_t exp32 = v.ui >> fp32FractionBits;\n",
      "      int32_t exp16 = exp32 - expAdjust;\n",
      "      // If exp16 == 0 (just into the denorm range), then significant should be shifted right 1.\n",
      "      // Smaller (so negative) exp16 values should result in greater right shifts.\n",
      "      uint32_t vshift = 1 - exp16;\n",
      "      uint32_t significand = fp32HiddenBit | (v.ui & fp32FractionMask);\n",
      "      v.ui = significand >> vshift;\n",
      "      v.ui += (v.ui & 0x3fff) != 0x1000 || (significand & 0x7ff) ? 0x1000 : 0;\n",
      "    } else if (v.si <= maxN) {\n",
      "      // Handle norms\n",
      "      v.ui += (v.ui & 0x3fff) != 0x1000 ? 0x1000 : 0;\n",
      "      v.ui -= expAdjust << fp32FractionBits;\n",
      "    } else if (v.si <= infN) {\n",
      "      v.si = infN;\n",
      "    } else if (v.si < nanN) {\n",
      "      v.si = nanN;\n",
      "    }\n",
      "\n",
      "    v.ui >>= shift;\n",
      "    return sign | (v.ui & 0x7fff);\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE float half2float(const uint16_t& value) const {\n",
      "    Bits v;\n",
      "    v.ui = value;\n",
      "    int32_t sign = v.si & signC;\n",
      "    v.si ^= sign;\n",
      "    sign <<= shiftSign;\n",
      "    v.si ^= ((v.si + minD) ^ v.si) & -(v.si > subC);\n",
      "    v.si ^= ((v.si + maxD) ^ v.si) & -(v.si > maxC);\n",
      "    Bits s;\n",
      "    s.si = mulC;\n",
      "    s.f *= v.si;\n",
      "    int32_t mask = -(norC > v.si);\n",
      "    v.si <<= shift;\n",
      "    v.si ^= (s.si ^ v.si) & mask;\n",
      "    v.si |= sign;\n",
      "    return v.f;\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE float half2float(\n",
      "    const volatile uint16_t& value) const volatile {\n",
      "    Bits v;\n",
      "    v.ui = value;\n",
      "    int32_t sign = v.si & signC;\n",
      "    v.si ^= sign;\n",
      "    sign <<= shiftSign;\n",
      "    v.si ^= ((v.si + minD) ^ v.si) & -(v.si > subC);\n",
      "    v.si ^= ((v.si + maxD) ^ v.si) & -(v.si > maxC);\n",
      "    Bits s;\n",
      "    s.si = mulC;\n",
      "    s.f *= v.si;\n",
      "    int32_t mask = -(norC > v.si);\n",
      "    v.si <<= shift;\n",
      "    v.si ^= (s.si ^ v.si) & mask;\n",
      "    v.si |= sign;\n",
      "    return v.f;\n",
      "  }\n",
      "\n",
      "  template<typename T>\n",
      "  TVM_XINLINE void constructor(const T& value) {\n",
      "    half_ = float2half(float(value));\n",
      "  }\n",
      "};\n",
      "\n",
      "TVM_HALF_OPERATOR(half, +)\n",
      "TVM_HALF_OPERATOR(half, -)\n",
      "TVM_HALF_OPERATOR(half, *)\n",
      "TVM_HALF_OPERATOR(half, /)\n",
      "TVM_HALF_OPERATOR(bool, >)\n",
      "TVM_HALF_OPERATOR(bool, <)\n",
      "TVM_HALF_OPERATOR(bool, >=)\n",
      "TVM_HALF_OPERATOR(bool, <=)\n",
      "\n",
      "TVM_XINLINE half __float2half_rn(const float a) {\n",
      "  return half(a);\n",
      "}\n",
      "#else\n",
      "#include <cuda_fp16.h>\n",
      "__device__ half max(half a, half b)\n",
      "{\n",
      "  return __hgt(__half(a), __half(b)) ? a : b;\n",
      "}\n",
      "__device__ half min(half a, half b)\n",
      "{\n",
      "  return __hlt(__half(a), __half(b)) ? a : b;\n",
      "}\n",
      "#endif\n",
      "\n",
      "\n",
      "// Pack two half values.\n",
      "static inline __device__ __host__ unsigned\n",
      "__pack_half2(const half x, const half y) {\n",
      "  unsigned v0 = *((unsigned short *)&x);\n",
      "  unsigned v1 = *((unsigned short *)&y);\n",
      "  return (v1 << 16) | v0;\n",
      "}\n",
      "\n",
      "#define CUDA_UNSUPPORTED_HALF_MATH_BINARY(HALF_MATH_NAME, FP32_MATH_NAME) \\\n",
      "static inline __device__ __host__ half HALF_MATH_NAME(half x, half y) {   \\\n",
      "  float tmp_x = __half2float(x);                                          \\\n",
      "  float tmp_y = __half2float(y);                                          \\\n",
      "  float result = FP32_MATH_NAME(tmp_x, tmp_y);                            \\\n",
      "  return __float2half(result);                                            \\\n",
      "}\n",
      "\n",
      "#define CUDA_UNSUPPORTED_HALF_MATH_UNARY(HALF_MATH_NAME, FP32_MATH_NAME) \\\n",
      "static inline __device__ __host__ half HALF_MATH_NAME(half x) {          \\\n",
      "  float tmp_x = __half2float(x);                                         \\\n",
      "  float result = FP32_MATH_NAME(tmp_x);                                  \\\n",
      "  return __float2half(result);                                           \\\n",
      "}\n",
      "\n",
      "// Some fp16 math functions are not supported in cuda_fp16.h,\n",
      "// so we define them here to make sure the generated CUDA code\n",
      "// is valid.\n",
      "#if defined(__CUDA_ARCH__)\n",
      "#if (__CUDA_ARCH__ >= 530)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_BINARY(hpow, powf)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(htanh, tanhf)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(htan, tanf)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(hatan, atanf)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(herf, erf)\n",
      "#else\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(hexp, exp)\n",
      "#endif\n",
      "#endif\n",
      "\n",
      "#undef CUDA_UNSUPPORTED_HALF_MATH_BINARY\n",
      "#undef CUDA_UNSUPPORTED_HALF_MATH_UNARY\n",
      "\n",
      "struct __align__(8) half4 {\n",
      "  __half x, y, z, w;\n",
      "  __host__ __device__ half4() : x(__half(0)), y(__half(0)), z(__half(0)), w(__half(0)) {}\n",
      "  __host__ __device__ half4(__half x, __half y, __half z, __half w) : x(x), y(y), z(z), w(w) {}\n",
      "\n",
      "};\n",
      "__host__ __device__ half4 make_half4(__half x, __half y, __half z, __half w) {\n",
      "    return half4(x, y, z, w);\n",
      "}\n",
      "#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 610)\n",
      "#include <sm_61_intrinsics.h>\n",
      "\n",
      "\n",
      "#if defined(__CUDACC_RTC__)\n",
      "#define __SM_61_INTRINSICS_DECL__ __device__\n",
      "#else /* !__CUDACC_RTC__ */\n",
      "#define __SM_61_INTRINSICS_DECL__ static __device__ __inline__\n",
      "#endif /* __CUDACC_RTC__ */\n",
      "\n",
      "#ifndef __CUDA_ARCH__\n",
      "#define __DEF_IF_HOST { }\n",
      "#else  /* !__CUDA_ARCH__ */\n",
      "#define __DEF_IF_HOST ;\n",
      "#endif /* __CUDA_ARCH__ */\n",
      "\n",
      "__SM_61_INTRINSICS_DECL__ int __dp4a(unsigned int srcA, int srcB, int c) __DEF_IF_HOST\n",
      "__SM_61_INTRINSICS_DECL__ int __dp4a(int srcA, unsigned int srcB, int c) __DEF_IF_HOST\n",
      "\n",
      "#undef __DEF_IF_HOST\n",
      "\n",
      "#if !defined(__CUDACC_RTC__) && defined(__CUDA_ARCH__)\n",
      "__SM_61_INTRINSICS_DECL__ int __dp4a(unsigned int srcA, int srcB, int c) {\n",
      "    int ret;\n",
      "    asm volatile (\"dp4a.u32.s32 %0, %1, %2, %3;\" : \"=r\"(ret) : \"r\"(srcA), \"r\"(srcB), \"r\"(c));\n",
      "    return ret;\n",
      "}\n",
      "\n",
      "__SM_61_INTRINSICS_DECL__ int __dp4a(int srcA, unsigned int srcB, int c) {\n",
      "    int ret;\n",
      "    asm volatile (\"dp4a.s32.u32 %0, %1, %2, %3;\" : \"=r\"(ret) : \"r\"(srcA), \"r\"(srcB), \"r\"(c));\n",
      "    return ret;\n",
      "}\n",
      "#endif /* !__CUDACC_RTC__ && defined(__CUDA_ARCH__) */\n",
      "\n",
      "#undef __SM_61_INTRINSICS_DECL__\n",
      "\n",
      "#endif\n",
      "__forceinline__ __device__ unsigned int\n",
      "cast_smem_ptr_to_int(const void* const smem_ptr)\n",
      "{\n",
      "  unsigned int smem_int;\n",
      "  asm volatile (\"{ .reg .u64 smem_int; cvta.to.shared.u64 smem_int, %1; cvt.u32.u64 %0, smem_int; }\"\n",
      "    : \"=r\"(smem_int) : \"l\"(smem_ptr));\n",
      "  return smem_int;\n",
      "}\n",
      "\n",
      "#if (((__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ >= 4)) || \\\n",
      "     (__CUDACC_VER_MAJOR__ > 11))\n",
      "#define TVM_ENABLE_L2_PREFETCH 1\n",
      "#else\n",
      "#define TVM_ENABLE_L2_PREFETCH 0\n",
      "#endif\n",
      "\n",
      "#ifdef _WIN32\n",
      "  using uint = unsigned int;\n",
      "  using uchar = unsigned char;\n",
      "  using ushort = unsigned short;\n",
      "  using int64_t = long long;\n",
      "  using uint64_t = unsigned long long;\n",
      "#else\n",
      "  #define uint unsigned int\n",
      "  #define uchar unsigned char\n",
      "  #define ushort unsigned short\n",
      "  #define int64_t long long\n",
      "  #define uint64_t unsigned long long\n",
      "#endif\n",
      "\n",
      "#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ == 800) \n",
      "#define TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST 1\n",
      "#else\n",
      "#define TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST 0\n",
      "#endif\n",
      "extern \"C\" __global__ void __launch_bounds__(128) matmul_n1024k1024_f16xi4_default_kernel(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m);\n",
      "extern \"C\" __global__ void __launch_bounds__(128) matmul_n1024k1024_f16xi4_default_kernel(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m) {\n",
      "  extern __shared__ uchar buf_dyn_shmem[];\n",
      "  half C_reindex_pad_shared_dyn_warp[128];\n",
      "  signed char B_local[4];\n",
      "  half B_decode_reindex_local[8];\n",
      "  half A_reindex_pad_shared_dyn_warp[32];\n",
      "  half B_decode_reindex_shared_dyn_warp[32];\n",
      "  for (int var = 0; var < 1; ++var) {\n",
      "    for (int ax1_0_3_init = 0; ax1_0_3_init < 4; ++ax1_0_3_init) {\n",
      "      for (int ax2_0_3_init = 0; ax2_0_3_init < 4; ++ax2_0_3_init) {\n",
      "        for (int i = 0; i < 8; ++i) {\n",
      "C_reindex_pad_shared_dyn_warp[((ax1_0_3_init * 32) + (ax2_0_3_init * 8)) + i] = 0.0;}\n",
      ";\n",
      "      }\n",
      "    }\n",
      "    for (int ax3_0_0 = 0; ax3_0_0 < 32; ++ax3_0_0) {\n",
      "      __syncthreads();\n",
      "      #pragma unroll\n",
      "      for (int ax0_ax1_ax2_fused_0 = 0; ax0_ax1_ax2_fused_0 < 8; ++ax0_ax1_ax2_fused_0) {\n",
      "        half4 condval;\n",
      "        if (((((((((int)blockIdx.y) * 128) + (ax0_ax1_ax2_fused_0 * 16)) + (((int)threadIdx.y) * 8)) + (((int)threadIdx.z) * 4)) + (((int)threadIdx.x) >> 3)) < m)) {\n",
      "          condval = *(half4*)(A + (((((((((int)blockIdx.y) * 131072) + (ax0_ax1_ax2_fused_0 * 16384)) + (((int)threadIdx.y) * 8192)) + (((int)threadIdx.z) * 4096)) + ((((int)threadIdx.x) >> 3) * 1024)) + (ax3_0_0 * 32)) + ((((int)threadIdx.x) & 7) * 4)));\n",
      "        } else {\n",
      "          condval = make_half4(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f));\n",
      "        }\n",
      "        *(half4*)(((half*)buf_dyn_shmem) + (((((((ax0_ax1_ax2_fused_0 * 512) + (((int)threadIdx.y) * 256)) + (((int)threadIdx.z) * 128)) + ((((int)threadIdx.x) >> 3) * 32)) + ((((((int)threadIdx.x) & 7) >> 1) ^ ((((int)threadIdx.z) * 2) + (((int)threadIdx.x) >> 4))) * 8)) + ((((int)threadIdx.x) & 1) * 4)) + 4096)) = condval;\n",
      "      }\n",
      "      for (int ax1_ax2_0_fused_0 = 0; ax1_ax2_0_fused_0 < 4; ++ax1_ax2_0_fused_0) {\n",
      "        *(int*)(B_local + 0) = *(int*)(B + (((((((((int)blockIdx.x) * 65536) + (ax1_ax2_0_fused_0 * 16384)) + (((int)threadIdx.y) * 8192)) + (((int)threadIdx.z) * 4096)) + ((((int)threadIdx.x) >> 2) * 512)) + (ax3_0_0 * 16)) + ((((int)threadIdx.x) & 3) * 4)));\n",
      "        decode_i4s_to_f16((&(B_local[0])), (&(B_decode_reindex_local[0])), 8);\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((((ax1_ax2_0_fused_0 * 1024) + (((int)threadIdx.y) * 512)) + (((int)threadIdx.z) * 256)) + ((((int)threadIdx.x) >> 2) * 32)) + (((((int)threadIdx.x) & 3) ^ (((int)threadIdx.x) >> 3)) * 8))) = *(uint4*)(B_decode_reindex_local + 0);\n",
      "      }\n",
      "      __syncthreads();\n",
      "      for (int ax3_0_1 = 0; ax3_0_1 < 2; ++ax3_0_1) {\n",
      "        for (int ax1_0 = 0; ax1_0 < 4; ++ax1_0) {\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[(((((((int)threadIdx.y) * 2048) + (ax1_0 * 512)) + ((((int)threadIdx.x) & 15) * 32)) + ((((ax3_0_1 * 2) + (((int)threadIdx.x) >> 4)) ^ ((((int)threadIdx.x) & 7) >> 1)) * 8)) + 4096)])) + 0)));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[(((((((int)threadIdx.y) * 2048) + (ax1_0 * 512)) + ((((int)threadIdx.x) & 15) * 32)) + ((((ax3_0_1 * 2) + (((int)threadIdx.x) >> 4)) ^ ((((int)threadIdx.x) & 7) >> 1)) * 8)) + 4096)])) + 0))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + (ax1_0 * 8)))[0]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + (ax1_0 * 8)))[1]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + (ax1_0 * 8)))[2]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + (ax1_0 * 8)))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "        }\n",
      "        for (int ax1_0_1 = 0; ax1_0_1 < 4; ++ax1_0_1) {\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[(((((((int)threadIdx.z) * 2048) + (ax1_0_1 * 512)) + ((((int)threadIdx.x) >> 4) * 256)) + ((((int)threadIdx.x) & 7) * 32)) + ((((ax3_0_1 * 2) + ((((int)threadIdx.x) & 15) >> 3)) ^ ((((int)threadIdx.x) & 7) >> 1)) * 8))])) + 0)));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[(((((((int)threadIdx.z) * 2048) + (ax1_0_1 * 512)) + ((((int)threadIdx.x) >> 4) * 256)) + ((((int)threadIdx.x) & 7) * 32)) + ((((ax3_0_1 * 2) + ((((int)threadIdx.x) & 15) >> 3)) ^ ((((int)threadIdx.x) & 7) >> 1)) * 8))])) + 0))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + (ax1_0_1 * 8)))[0]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + (ax1_0_1 * 8)))[1]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + (ax1_0_1 * 8)))[2]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + (ax1_0_1 * 8)))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "        }\n",
      "        for (int ax1_0_3 = 0; ax1_0_3 < 4; ++ax1_0_3) {\n",
      "          for (int ax2_0_3 = 0; ax2_0_3 < 4; ++ax2_0_3) {\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + ((ax1_0_3 * 32) + (ax2_0_3 * 8))))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + ((ax1_0_3 * 32) + (ax2_0_3 * 8))))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + (ax1_0_3 * 8)))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + (ax1_0_3 * 8)))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + (ax1_0_3 * 8)))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + (ax1_0_3 * 8)))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + (ax2_0_3 * 8)))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + (ax2_0_3 * 8)))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + ((ax1_0_3 * 32) + (ax2_0_3 * 8))))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + ((ax1_0_3 * 32) + (ax2_0_3 * 8))))[1]));\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + (((ax1_0_3 * 32) + (ax2_0_3 * 8)) + 4)))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + (((ax1_0_3 * 32) + (ax2_0_3 * 8)) + 4)))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + (ax1_0_3 * 8)))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + (ax1_0_3 * 8)))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + (ax1_0_3 * 8)))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + (ax1_0_3 * 8)))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + ((ax2_0_3 * 8) + 4)))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + ((ax2_0_3 * 8) + 4)))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + (((ax1_0_3 * 32) + (ax2_0_3 * 8)) + 4)))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + (((ax1_0_3 * 32) + (ax2_0_3 * 8)) + 4)))[1]));\n",
      "  }\n",
      "          }\n",
      "        }\n",
      "      }\n",
      "    }\n",
      "    for (int ax0 = 0; ax0 < 4; ++ax0) {\n",
      "      __syncthreads();\n",
      "      for (int ax1 = 0; ax1 < 4; ++ax1) {\n",
      "        for (int local_id = 0; local_id < 8; local_id+=2) {\n",
      "*((uint *)&(&(((half*)buf_dyn_shmem)[(((((int)threadIdx.y) * 8192) + (((int)threadIdx.z) * 1024)) + (ax1 * 256))]))[((((((local_id % 4) / 2) * 8) + (threadIdx.x / 4)) * 16) + ((((local_id / 4) * 8) + ((threadIdx.x % 4) * 2)) + (local_id % 2)))]) = *((uint *)&C_reindex_pad_shared_dyn_warp[((ax0 * 32) + (ax1 * 8)) + local_id]);\n",
      "}\n",
      ";\n",
      "      }\n",
      "      __syncthreads();\n",
      "      #pragma unroll\n",
      "      for (int ax0_ax1_ax2_ax3_ax4_fused_0 = 0; ax0_ax1_ax2_ax3_ax4_fused_0 < 4; ++ax0_ax1_ax2_ax3_ax4_fused_0) {\n",
      "        if (((((((int)blockIdx.y) * 128) + (((int)threadIdx.y) * 64)) + (ax0 * 16)) + (((int)threadIdx.x) >> 1)) < m) {\n",
      "          *(uint4*)(C + ((((((((((int)blockIdx.y) * 131072) + (((int)threadIdx.y) * 65536)) + (ax0 * 16384)) + ((((int)threadIdx.x) >> 1) * 1024)) + (((int)blockIdx.x) * 128)) + (((int)threadIdx.z) * 64)) + (ax0_ax1_ax2_ax3_ax4_fused_0 * 16)) + ((((int)threadIdx.x) & 1) * 8))) = *(uint4*)(((half*)buf_dyn_shmem) + ((((((int)threadIdx.y) * 8192) + (((int)threadIdx.z) * 1024)) + (ax0_ax1_ax2_ax3_ax4_fused_0 * 256)) + (((int)threadIdx.x) * 8)));\n",
      "        }\n",
      "      }\n",
      "    }\n",
      "  }\n",
      "}\n",
      "\n",
      "\n",
      "extern \"C\" void init() {\n",
      "    \n",
      "    cudaFuncSetAttribute(matmul_n1024k1024_f16xi4_default_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 20480);\n",
      "\n",
      "}\n",
      "\n",
      "extern \"C\" void call(half* __restrict__ A, int8_t* __restrict__ B, half* __restrict__ C, int m, cudaStream_t stream=cudaStreamDefault) {\n",
      "if (m == 0) return; \n",
      "\t\tmatmul_n1024k1024_f16xi4_default_kernel<<<dim3(8, (m + 127) / 128, 1), dim3(32, 2, 2), 20480, stream>>>(A, B, C, m);\n",
      "}\n",
      "\n"
     ]
    }
   ],
   "source": [
    "matmul_config = bitblas.MatmulConfig(\n",
    "    M=[1, 16, 32, 64, 128, 256],  # M dimension, default value is from 1 to 1024\n",
    "    N=1024,  # N dimension\n",
    "    K=1024,  # K dimension\n",
    "    A_dtype=\"float16\",  # activation A dtype\n",
    "    W_dtype=\"int4\",  # weight W dtype\n",
    ")\n",
    "\n",
    "matmul = bitblas.Matmul(config=matmul_config, enable_tuning=False)\n",
    "print(matmul.get_source())"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "id": "d701fa2e-ef3b-49c6-a54a-42c8f5219296",
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "2024-10-25 08:27:16 [BitBLAS:INFO]: Start fast tuning with dynamic range\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Cannot find the appropriate index map for tensorcore\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Apply config {'block': [2], 'thread': [2], 'rstep': [1024], 'reduce_thread': [64], 'vectorize': {'A': 8, 'B_decode': 8}}\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Apply config {'block': [1], 'thread': [1], 'rstep': [1024], 'reduce_thread': [128], 'vectorize': {'A': 8, 'B_decode': 8}}\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Apply config {'block': [4], 'thread': [4], 'rstep': [1024], 'reduce_thread': [32], 'vectorize': {'A': 8, 'B_decode': 8}}\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Apply config {'block': [8], 'thread': [8], 'rstep': [1024], 'reduce_thread': [16], 'vectorize': {'A': 8, 'B_decode': 8}}\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Apply config {'block': [16], 'thread': [16], 'rstep': [512], 'reduce_thread': [8], 'vectorize': {'A': 4, 'B_decode': 8}}\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Apply config {'block': [32], 'thread': [32], 'rstep': [256], 'reduce_thread': [4], 'vectorize': {'A': 2, 'B_decode': 8}}\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Apply config {'block': [64], 'thread': [64], 'rstep': [128], 'reduce_thread': [2], 'vectorize': {'B_decode': 8}}\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Apply config {'block': [128], 'thread': [128], 'rstep': [128], 'vectorize': {'B_decode': 8}}\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Warning: block config [128] is not valid for matmul, skip.\n",
      "2024-10-25 08:27:16 [BitBLAS:DEBUG]: Warning: block config [128] is not valid for matmul, skip.\n",
      "2024-10-25 08:27:41 [BitBLAS:DEBUG]: LocalBuilder: An exception occurred Traceback (most recent call last):\n",
      "  File \"/root/BitBLAS/3rdparty/tvm/python/tvm/exec/popen_worker.py\", line 87, in main\n",
      "    result = fn(*args, **kwargs)\n",
      "  File \"/root/BitBLAS/bitblas/base/utils.py\", line 257, in _build\n",
      "    rt_mod = tvm.build(mod, ta\t...\tm::codegen::CodeGenCUDA::VisitExpr_(tvm::tir::RampNode const*, std::ostream&)\n",
      "  File \"/root/BitBLAS/3rdparty/tvm/src/target/source/codegen_cuda.cc\", line 1226\n",
      "ValueError: Check failed: lanes <= 4 (8 vs. 4) : Ramp of more than 4 lanes is not allowed.\n",
      "\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Evaluation with config {'block': [2], 'thread': [2], 'rstep': [1024], 'reduce_thread': [64], 'vectorize': {'A': 8, 'B_decode': 8}}\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Time cost of this config: 0.007 ms\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Evaluation with config {'block': [1], 'thread': [1], 'rstep': [1024], 'reduce_thread': [128], 'vectorize': {'A': 8, 'B_decode': 8}}\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Time cost of this config: 0.006 ms\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Evaluation with config {'block': [4], 'thread': [4], 'rstep': [1024], 'reduce_thread': [32], 'vectorize': {'A': 8, 'B_decode': 8}}\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Time cost of this config: 0.005 ms\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Evaluation with config {'block': [8], 'thread': [8], 'rstep': [1024], 'reduce_thread': [16], 'vectorize': {'A': 8, 'B_decode': 8}}\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Time cost of this config: 0.005 ms\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Evaluation with config {'block': [16], 'thread': [16], 'rstep': [512], 'reduce_thread': [8], 'vectorize': {'A': 4, 'B_decode': 8}}\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Time cost of this config: 0.005 ms\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Evaluation with config {'block': [32], 'thread': [32], 'rstep': [256], 'reduce_thread': [4], 'vectorize': {'A': 2, 'B_decode': 8}}\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Time cost of this config: 0.007 ms\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Evaluation with config {'block': [64], 'thread': [64], 'rstep': [128], 'reduce_thread': [2], 'vectorize': {'B_decode': 8}}\n",
      "2024-10-25 08:27:42 [BitBLAS:INFO]: Time cost of this config: 0.005 ms\n",
      "2024-10-25 08:27:43 [BitBLAS:DEBUG]: Apply config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:27:43 [BitBLAS:DEBUG]: Apply config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:27:43 [BitBLAS:DEBUG]: Apply config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:27:43 [BitBLAS:DEBUG]: Apply config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:27:43 [BitBLAS:DEBUG]: Apply config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:27:43 [BitBLAS:DEBUG]: Apply config {'block': [16, 512], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Evaluation with config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Time cost of this config: 0.009 ms\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Evaluation with config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Time cost of this config: 0.011 ms\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Evaluation with config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Time cost of this config: 0.009 ms\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Evaluation with config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Time cost of this config: 0.007 ms\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Evaluation with config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Time cost of this config: 0.049 ms\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Evaluation with config {'block': [16, 512], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:27 [BitBLAS:INFO]: Time cost of this config: 0.024 ms\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [32, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [32, 32], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [32, 64], 'warp': [16, 32], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [32, 128], 'warp': [16, 64], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [32, 256], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:28 [BitBLAS:DEBUG]: Apply config {'block': [16, 512], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:28:29 [BitBLAS:DEBUG]: Apply config {'block': [32, 512], 'warp': [16, 256], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.008 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [32, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.009 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.007 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [32, 32], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.009 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [32, 64], 'warp': [16, 32], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.024 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.008 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [32, 128], 'warp': [16, 64], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.015 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [32, 256], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.030 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [16, 512], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.048 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Evaluation with config {'block': [32, 512], 'warp': [16, 256], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:INFO]: Time cost of this config: 0.040 ms\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [32, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [32, 32], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [32, 64], 'warp': [16, 32], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [64, 16], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [64, 32], 'warp': [32, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:16 [BitBLAS:DEBUG]: Apply config {'block': [32, 128], 'warp': [16, 64], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:17 [BitBLAS:DEBUG]: Apply config {'block': [64, 64], 'warp': [32, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:17 [BitBLAS:DEBUG]: Apply config {'block': [64, 128], 'warp': [32, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:17 [BitBLAS:DEBUG]: Apply config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:17 [BitBLAS:DEBUG]: Apply config {'block': [32, 256], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:17 [BitBLAS:DEBUG]: Apply config {'block': [64, 256], 'warp': [32, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:17 [BitBLAS:DEBUG]: Apply config {'block': [16, 512], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:17 [BitBLAS:DEBUG]: Apply config {'block': [32, 512], 'warp': [16, 256], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:29:17 [BitBLAS:DEBUG]: Apply config {'block': [64, 512], 'warp': [32, 256], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.009 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [32, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.008 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.009 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [32, 32], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.009 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [32, 64], 'warp': [16, 32], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [64, 16], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.015 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [64, 32], 'warp': [32, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [32, 128], 'warp': [16, 64], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.011 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [64, 64], 'warp': [32, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.015 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [64, 128], 'warp': [32, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.021 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.048 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [32, 256], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.030 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [64, 256], 'warp': [32, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.024 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [16, 512], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.031 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [32, 512], 'warp': [16, 256], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.040 ms\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Evaluation with config {'block': [64, 512], 'warp': [32, 256], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:04 [BitBLAS:INFO]: Time cost of this config: 0.055 ms\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [32, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [32, 32], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [32, 64], 'warp': [16, 32], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [64, 16], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [64, 32], 'warp': [32, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:05 [BitBLAS:DEBUG]: Apply config {'block': [32, 128], 'warp': [16, 64], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [64, 64], 'warp': [32, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [64, 128], 'warp': [32, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [128, 16], 'warp': [32, 16], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [128, 32], 'warp': [64, 16], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [32, 256], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [128, 64], 'warp': [64, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [64, 256], 'warp': [32, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [128, 128], 'warp': [64, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:06 [BitBLAS:DEBUG]: Apply config {'block': [16, 512], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [32, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.011 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [32, 32], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.011 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [32, 64], 'warp': [16, 32], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.008 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [64, 16], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.015 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [64, 32], 'warp': [32, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [32, 128], 'warp': [16, 64], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.008 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [64, 64], 'warp': [32, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.015 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [64, 128], 'warp': [32, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.021 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [128, 16], 'warp': [32, 16], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.030 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.024 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [128, 32], 'warp': [64, 16], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.020 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [32, 256], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.015 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [128, 64], 'warp': [64, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.018 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [64, 256], 'warp': [32, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.048 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [128, 128], 'warp': [64, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.030 ms\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Evaluation with config {'block': [16, 512], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:54 [BitBLAS:INFO]: Time cost of this config: 0.029 ms\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [32, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [32, 32], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [32, 64], 'warp': [16, 32], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [64, 16], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [64, 32], 'warp': [32, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:55 [BitBLAS:DEBUG]: Apply config {'block': [32, 128], 'warp': [16, 64], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [64, 64], 'warp': [32, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [64, 128], 'warp': [32, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [128, 16], 'warp': [32, 16], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [128, 32], 'warp': [64, 16], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [32, 256], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [128, 64], 'warp': [64, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [64, 256], 'warp': [32, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [128, 128], 'warp': [64, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:30:56 [BitBLAS:DEBUG]: Apply config {'block': [256, 16], 'warp': [64, 16], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 4}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [16, 32], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [32, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.011 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [16, 16], 'warp': [16, 16], 'rstep': [256], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.013 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [16, 64], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.014 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [32, 32], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.015 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [32, 64], 'warp': [16, 32], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.010 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [64, 16], 'warp': [16, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.011 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [16, 128], 'warp': [16, 32], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.013 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [64, 32], 'warp': [32, 16], 'rstep': [128], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.016 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [32, 128], 'warp': [16, 64], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.012 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [64, 64], 'warp': [32, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.021 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [64, 128], 'warp': [32, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.015 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [128, 16], 'warp': [32, 16], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.016 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [16, 256], 'warp': [16, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 4, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.024 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [128, 32], 'warp': [64, 16], 'rstep': [64], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.018 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [32, 256], 'warp': [16, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.030 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [128, 64], 'warp': [64, 32], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.021 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [64, 256], 'warp': [32, 128], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.030 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [128, 128], 'warp': [64, 64], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 8}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.030 ms\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Evaluation with config {'block': [256, 16], 'warp': [64, 16], 'rstep': [32], 'use_tc': True, 'vectorize': {'A_reindex': 8, 'B_decode_reindex': 4}}\n",
      "2024-10-25 08:31:43 [BitBLAS:INFO]: Time cost of this config: 0.031 ms\n",
      "\n",
      "template <typename T1, typename T2, bool isSigned = false>\n",
      "__device__ void decode_i4b_to_f16(T1 *_i4s, T2 *B_local_decode, const int N = 8)\n",
      "{\n",
      "    uint *h = reinterpret_cast<uint *>(B_local_decode);\n",
      "\n",
      "    static constexpr uint immLut = (0xf0 & 0xcc) | 0xaa;\n",
      "    static constexpr uint BOTTOM_MASK = 0x000f000f;\n",
      "    static constexpr uint FP16_TOP_MAGIC_NUM = 0x64006400;\n",
      "    static constexpr uint MEDIAN_NUM = isSigned ? 0x64086408 : 0x64006400;\n",
      "    uint const i4s = *reinterpret_cast<uint *>(_i4s);\n",
      "#pragma unroll\n",
      "    for (int i = 0; i < (N / 2); i++)\n",
      "    {\n",
      "\n",
      "        asm volatile(\"lop3.b32 %0, %1, %2, %3, %4;\\n\"\n",
      "                     : \"=r\"(h[i])\n",
      "                     : \"r\"(i4s >> (4 * i)), \"n\"(BOTTOM_MASK), \"n\"(FP16_TOP_MAGIC_NUM), \"n\"(immLut));\n",
      "        asm volatile(\"sub.f16x2 %0, %1, %2;\\n\" : \"=r\"(h[i]) : \"r\"(h[i]), \"r\"(MEDIAN_NUM));\n",
      "    }\n",
      "}\n",
      "\n",
      "template <typename T1, typename T2>\n",
      "__device__ void decode_i4s_to_f16(T1 *_i4s, T2 *B_local_decode, const int N = 8)\n",
      "{\n",
      "    decode_i4b_to_f16<T1, T2, true>(_i4s, B_local_decode, N);\n",
      "}\n",
      "\n",
      "template <typename T1, typename T2>\n",
      "__device__ void decode_i4u_to_f16(T1 *_i4u, T2 *B_local_decode, const int N = 8)\n",
      "{\n",
      "    decode_i4b_to_f16<T1, T2, false>(_i4u, B_local_decode, N);\n",
      "}\n",
      "#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 530)\n",
      "\n",
      "typedef unsigned short uint16_t;\n",
      "typedef unsigned char uint8_t;\n",
      "typedef signed char int8_t;\n",
      "typedef int int32_t;\n",
      "typedef unsigned long long uint64_t;\n",
      "typedef unsigned int uint32_t;\n",
      "\n",
      "#define TVM_FORCE_INLINE inline __attribute__((always_inline))\n",
      "#define TVM_XINLINE TVM_FORCE_INLINE __device__ __host__\n",
      "#define TVM_ALIGNED(x) __attribute__ ((aligned(x)))\n",
      "#define TVM_HALF_OPERATOR(RTYPE, OP)                              \\\n",
      "  TVM_XINLINE RTYPE operator OP (half a, half b) {                \\\n",
      "    return RTYPE(float(a) OP float(b));                           \\\n",
      "  }                                                               \\\n",
      "  template<typename T>                                            \\\n",
      "  TVM_XINLINE RTYPE operator OP (half a, T b) {                   \\\n",
      "    return RTYPE(float(a) OP float(b));                           \\\n",
      "  }                                                               \\\n",
      "  template<typename T>                                            \\\n",
      "  TVM_XINLINE RTYPE operator OP (T a, half b) {                   \\\n",
      "    return RTYPE(float(a) OP float(b));                           \\\n",
      "  }\n",
      "\n",
      "#define TVM_HALF_ASSIGNOP(AOP, OP)                                \\\n",
      "  template<typename T>                                            \\\n",
      "  TVM_XINLINE half operator AOP (const T& a) {                    \\\n",
      "    return *this = half(float(*this) OP float(a));                \\\n",
      "  }                                                               \\\n",
      "  template<typename T>                                            \\\n",
      "  TVM_XINLINE half operator AOP (const volatile T& a) volatile {  \\\n",
      "    return *this = half(float(*this) OP float(a));                \\\n",
      "  }\n",
      "\n",
      "class TVM_ALIGNED(2) half {\n",
      " public:\n",
      "  uint16_t half_;\n",
      "\n",
      "  static TVM_XINLINE half Binary(uint16_t value) {\n",
      "    half res;\n",
      "    res.half_ = value;\n",
      "    return res;\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE half() {}\n",
      "\n",
      "  TVM_XINLINE half(const float& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const double& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const int8_t& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const uint8_t& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const int32_t& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const uint32_t& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const long long& value) { constructor(value); }\n",
      "  TVM_XINLINE explicit half(const uint64_t& value) { constructor(value); }\n",
      "\n",
      "  TVM_XINLINE operator float() const {                          \\\n",
      "    return float(half2float(half_));                            \\\n",
      "  }                                                             \\\n",
      "  TVM_XINLINE operator float() const volatile {                 \\\n",
      "    return float(half2float(half_));                            \\\n",
      "  }\n",
      "\n",
      "\n",
      "  TVM_HALF_ASSIGNOP(+=, +)\n",
      "  TVM_HALF_ASSIGNOP(-=, -)\n",
      "  TVM_HALF_ASSIGNOP(*=, *)\n",
      "  TVM_HALF_ASSIGNOP(/=, /)\n",
      "\n",
      "  TVM_XINLINE half operator+() {\n",
      "    return *this;\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE half operator-() {\n",
      "    return half(-float(*this));\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE half operator=(const half& a) {\n",
      "    half_ = a.half_;\n",
      "    return a;\n",
      "  }\n",
      "\n",
      "  template<typename T>\n",
      "  TVM_XINLINE half operator=(const T& a) {\n",
      "    return *this = half(a);\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE half operator=(const half& a) volatile {\n",
      "    half_ = a.half_;\n",
      "    return a;\n",
      "  }\n",
      "\n",
      "  template<typename T>\n",
      "  TVM_XINLINE half operator=(const T& a) volatile {\n",
      "    return *this = half(a);\n",
      "  }\n",
      "\n",
      " private:\n",
      "  union Bits {\n",
      "    float f;\n",
      "    int32_t si;\n",
      "    uint32_t ui;\n",
      "  };\n",
      "\n",
      "  static int const fp16FractionBits = 10;\n",
      "  static int const fp32FractionBits = 23;\n",
      "  static int32_t const fp32FractionMask = ~(~0u << fp32FractionBits);   // == 0x7fffff\n",
      "  static int32_t const fp32HiddenBit = 1 << fp32FractionBits;   // == 0x800000\n",
      "  static int const shift = fp32FractionBits - fp16FractionBits;   // == 13\n",
      "  static int const shiftSign = 16;\n",
      "  static int32_t const expAdjust = 127 - 15;   // exp32-127 = exp16-15, so exp16 = exp32 - (127-15)\n",
      "\n",
      "  static int32_t const infN = 0x7F800000;   // flt32 infinity\n",
      "  static int32_t const maxN = 0x477FFFFF;   // max flt32 that's a flt16 normal after >> by shift\n",
      "  static int32_t const minN = 0x38800000;   // min flt16 normal as a flt32\n",
      "  static int32_t const maxZ = 0x33000000;   // max fp32 number that's still rounded to zero in fp16\n",
      "  static int32_t const signN = 0x80000000;  // flt32 sign bit\n",
      "\n",
      "  static int32_t const infC = infN >> shift;\n",
      "  static int32_t const nanN = (infC + 1) << shift;   // minimum flt16 nan as a flt32\n",
      "  static int32_t const maxC = maxN >> shift;\n",
      "  static int32_t const minC = minN >> shift;\n",
      "  static int32_t const signC = signN >> shiftSign;  // flt16 sign bit\n",
      "\n",
      "  static int32_t const mulN = 0x52000000;  // (1 << 23) / minN\n",
      "  static int32_t const mulC = 0x33800000;  // minN / (1 << (23 - shift))\n",
      "\n",
      "  static int32_t const subC = 0x003FF;  // max flt32 subnormal down shifted\n",
      "  static int32_t const norC = 0x00400;  // min flt32 normal down shifted\n",
      "\n",
      "  static int32_t const maxD = infC - maxC - 1;\n",
      "  static int32_t const minD = minC - subC - 1;\n",
      "\n",
      "  TVM_XINLINE uint16_t float2half(const float& value) const {\n",
      "    Bits v;\n",
      "    v.f = value;\n",
      "    uint32_t sign = v.si & signN;    // grab sign bit\n",
      "    v.si ^= sign;                    // clear sign bit from v\n",
      "    sign >>= shiftSign;              // logical shift sign to fp16 position\n",
      "\n",
      "    if (v.si <= maxZ) {\n",
      "      // Handle eventual zeros here to ensure\n",
      "      // vshift will not exceed 32 below.\n",
      "      v.ui = 0;\n",
      "    } else if (v.si < minN) {\n",
      "      // Handle denorms\n",
      "      uint32_t exp32 = v.ui >> fp32FractionBits;\n",
      "      int32_t exp16 = exp32 - expAdjust;\n",
      "      // If exp16 == 0 (just into the denorm range), then significant should be shifted right 1.\n",
      "      // Smaller (so negative) exp16 values should result in greater right shifts.\n",
      "      uint32_t vshift = 1 - exp16;\n",
      "      uint32_t significand = fp32HiddenBit | (v.ui & fp32FractionMask);\n",
      "      v.ui = significand >> vshift;\n",
      "      v.ui += (v.ui & 0x3fff) != 0x1000 || (significand & 0x7ff) ? 0x1000 : 0;\n",
      "    } else if (v.si <= maxN) {\n",
      "      // Handle norms\n",
      "      v.ui += (v.ui & 0x3fff) != 0x1000 ? 0x1000 : 0;\n",
      "      v.ui -= expAdjust << fp32FractionBits;\n",
      "    } else if (v.si <= infN) {\n",
      "      v.si = infN;\n",
      "    } else if (v.si < nanN) {\n",
      "      v.si = nanN;\n",
      "    }\n",
      "\n",
      "    v.ui >>= shift;\n",
      "    return sign | (v.ui & 0x7fff);\n",
      "  }\n",
      "\n",
      "  // Same as above routine, except for addition of volatile keyword\n",
      "  TVM_XINLINE uint16_t float2half(\n",
      "    const volatile float& value) const volatile {\n",
      "    Bits v;\n",
      "    v.f = value;\n",
      "    uint32_t sign = v.si & signN;    // grab sign bit\n",
      "    v.si ^= sign;                    // clear sign bit from v\n",
      "    sign >>= shiftSign;              // logical shift sign to fp16 position\n",
      "\n",
      "    if (v.si <= maxZ) {\n",
      "      // Handle eventual zeros here to ensure\n",
      "      // vshift will not exceed 32 below.\n",
      "      v.ui = 0;\n",
      "    } else if (v.si < minN) {\n",
      "      // Handle denorms\n",
      "      uint32_t exp32 = v.ui >> fp32FractionBits;\n",
      "      int32_t exp16 = exp32 - expAdjust;\n",
      "      // If exp16 == 0 (just into the denorm range), then significant should be shifted right 1.\n",
      "      // Smaller (so negative) exp16 values should result in greater right shifts.\n",
      "      uint32_t vshift = 1 - exp16;\n",
      "      uint32_t significand = fp32HiddenBit | (v.ui & fp32FractionMask);\n",
      "      v.ui = significand >> vshift;\n",
      "      v.ui += (v.ui & 0x3fff) != 0x1000 || (significand & 0x7ff) ? 0x1000 : 0;\n",
      "    } else if (v.si <= maxN) {\n",
      "      // Handle norms\n",
      "      v.ui += (v.ui & 0x3fff) != 0x1000 ? 0x1000 : 0;\n",
      "      v.ui -= expAdjust << fp32FractionBits;\n",
      "    } else if (v.si <= infN) {\n",
      "      v.si = infN;\n",
      "    } else if (v.si < nanN) {\n",
      "      v.si = nanN;\n",
      "    }\n",
      "\n",
      "    v.ui >>= shift;\n",
      "    return sign | (v.ui & 0x7fff);\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE float half2float(const uint16_t& value) const {\n",
      "    Bits v;\n",
      "    v.ui = value;\n",
      "    int32_t sign = v.si & signC;\n",
      "    v.si ^= sign;\n",
      "    sign <<= shiftSign;\n",
      "    v.si ^= ((v.si + minD) ^ v.si) & -(v.si > subC);\n",
      "    v.si ^= ((v.si + maxD) ^ v.si) & -(v.si > maxC);\n",
      "    Bits s;\n",
      "    s.si = mulC;\n",
      "    s.f *= v.si;\n",
      "    int32_t mask = -(norC > v.si);\n",
      "    v.si <<= shift;\n",
      "    v.si ^= (s.si ^ v.si) & mask;\n",
      "    v.si |= sign;\n",
      "    return v.f;\n",
      "  }\n",
      "\n",
      "  TVM_XINLINE float half2float(\n",
      "    const volatile uint16_t& value) const volatile {\n",
      "    Bits v;\n",
      "    v.ui = value;\n",
      "    int32_t sign = v.si & signC;\n",
      "    v.si ^= sign;\n",
      "    sign <<= shiftSign;\n",
      "    v.si ^= ((v.si + minD) ^ v.si) & -(v.si > subC);\n",
      "    v.si ^= ((v.si + maxD) ^ v.si) & -(v.si > maxC);\n",
      "    Bits s;\n",
      "    s.si = mulC;\n",
      "    s.f *= v.si;\n",
      "    int32_t mask = -(norC > v.si);\n",
      "    v.si <<= shift;\n",
      "    v.si ^= (s.si ^ v.si) & mask;\n",
      "    v.si |= sign;\n",
      "    return v.f;\n",
      "  }\n",
      "\n",
      "  template<typename T>\n",
      "  TVM_XINLINE void constructor(const T& value) {\n",
      "    half_ = float2half(float(value));\n",
      "  }\n",
      "};\n",
      "\n",
      "TVM_HALF_OPERATOR(half, +)\n",
      "TVM_HALF_OPERATOR(half, -)\n",
      "TVM_HALF_OPERATOR(half, *)\n",
      "TVM_HALF_OPERATOR(half, /)\n",
      "TVM_HALF_OPERATOR(bool, >)\n",
      "TVM_HALF_OPERATOR(bool, <)\n",
      "TVM_HALF_OPERATOR(bool, >=)\n",
      "TVM_HALF_OPERATOR(bool, <=)\n",
      "\n",
      "TVM_XINLINE half __float2half_rn(const float a) {\n",
      "  return half(a);\n",
      "}\n",
      "#else\n",
      "#include <cuda_fp16.h>\n",
      "__device__ half max(half a, half b)\n",
      "{\n",
      "  return __hgt(__half(a), __half(b)) ? a : b;\n",
      "}\n",
      "__device__ half min(half a, half b)\n",
      "{\n",
      "  return __hlt(__half(a), __half(b)) ? a : b;\n",
      "}\n",
      "#endif\n",
      "\n",
      "\n",
      "// Pack two half values.\n",
      "static inline __device__ __host__ unsigned\n",
      "__pack_half2(const half x, const half y) {\n",
      "  unsigned v0 = *((unsigned short *)&x);\n",
      "  unsigned v1 = *((unsigned short *)&y);\n",
      "  return (v1 << 16) | v0;\n",
      "}\n",
      "\n",
      "#define CUDA_UNSUPPORTED_HALF_MATH_BINARY(HALF_MATH_NAME, FP32_MATH_NAME) \\\n",
      "static inline __device__ __host__ half HALF_MATH_NAME(half x, half y) {   \\\n",
      "  float tmp_x = __half2float(x);                                          \\\n",
      "  float tmp_y = __half2float(y);                                          \\\n",
      "  float result = FP32_MATH_NAME(tmp_x, tmp_y);                            \\\n",
      "  return __float2half(result);                                            \\\n",
      "}\n",
      "\n",
      "#define CUDA_UNSUPPORTED_HALF_MATH_UNARY(HALF_MATH_NAME, FP32_MATH_NAME) \\\n",
      "static inline __device__ __host__ half HALF_MATH_NAME(half x) {          \\\n",
      "  float tmp_x = __half2float(x);                                         \\\n",
      "  float result = FP32_MATH_NAME(tmp_x);                                  \\\n",
      "  return __float2half(result);                                           \\\n",
      "}\n",
      "\n",
      "// Some fp16 math functions are not supported in cuda_fp16.h,\n",
      "// so we define them here to make sure the generated CUDA code\n",
      "// is valid.\n",
      "#if defined(__CUDA_ARCH__)\n",
      "#if (__CUDA_ARCH__ >= 530)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_BINARY(hpow, powf)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(htanh, tanhf)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(htan, tanf)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(hatan, atanf)\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(herf, erf)\n",
      "#else\n",
      "CUDA_UNSUPPORTED_HALF_MATH_UNARY(hexp, exp)\n",
      "#endif\n",
      "#endif\n",
      "\n",
      "#undef CUDA_UNSUPPORTED_HALF_MATH_BINARY\n",
      "#undef CUDA_UNSUPPORTED_HALF_MATH_UNARY\n",
      "\n",
      "struct __align__(8) half4 {\n",
      "  __half x, y, z, w;\n",
      "  __host__ __device__ half4() : x(__half(0)), y(__half(0)), z(__half(0)), w(__half(0)) {}\n",
      "  __host__ __device__ half4(__half x, __half y, __half z, __half w) : x(x), y(y), z(z), w(w) {}\n",
      "\n",
      "};\n",
      "__host__ __device__ half4 make_half4(__half x, __half y, __half z, __half w) {\n",
      "    return half4(x, y, z, w);\n",
      "}\n",
      "\n",
      "#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 700)\n",
      "#define __shfl_sync(mask, var, lane, width) \\\n",
      "        __shfl((var), (lane), (width))\n",
      "\n",
      "#define __shfl_down_sync(mask, var, offset, width) \\\n",
      "        __shfl_down((var), (offset), (width))\n",
      "\n",
      "#define __shfl_up_sync(mask, var, offset, width) \\\n",
      "        __shfl_up((var), (offset), (width))\n",
      "#endif\n",
      "\n",
      "#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 610)\n",
      "#include <sm_61_intrinsics.h>\n",
      "\n",
      "\n",
      "#if defined(__CUDACC_RTC__)\n",
      "#define __SM_61_INTRINSICS_DECL__ __device__\n",
      "#else /* !__CUDACC_RTC__ */\n",
      "#define __SM_61_INTRINSICS_DECL__ static __device__ __inline__\n",
      "#endif /* __CUDACC_RTC__ */\n",
      "\n",
      "#ifndef __CUDA_ARCH__\n",
      "#define __DEF_IF_HOST { }\n",
      "#else  /* !__CUDA_ARCH__ */\n",
      "#define __DEF_IF_HOST ;\n",
      "#endif /* __CUDA_ARCH__ */\n",
      "\n",
      "__SM_61_INTRINSICS_DECL__ int __dp4a(unsigned int srcA, int srcB, int c) __DEF_IF_HOST\n",
      "__SM_61_INTRINSICS_DECL__ int __dp4a(int srcA, unsigned int srcB, int c) __DEF_IF_HOST\n",
      "\n",
      "#undef __DEF_IF_HOST\n",
      "\n",
      "#if !defined(__CUDACC_RTC__) && defined(__CUDA_ARCH__)\n",
      "__SM_61_INTRINSICS_DECL__ int __dp4a(unsigned int srcA, int srcB, int c) {\n",
      "    int ret;\n",
      "    asm volatile (\"dp4a.u32.s32 %0, %1, %2, %3;\" : \"=r\"(ret) : \"r\"(srcA), \"r\"(srcB), \"r\"(c));\n",
      "    return ret;\n",
      "}\n",
      "\n",
      "__SM_61_INTRINSICS_DECL__ int __dp4a(int srcA, unsigned int srcB, int c) {\n",
      "    int ret;\n",
      "    asm volatile (\"dp4a.s32.u32 %0, %1, %2, %3;\" : \"=r\"(ret) : \"r\"(srcA), \"r\"(srcB), \"r\"(c));\n",
      "    return ret;\n",
      "}\n",
      "#endif /* !__CUDACC_RTC__ && defined(__CUDA_ARCH__) */\n",
      "\n",
      "#undef __SM_61_INTRINSICS_DECL__\n",
      "\n",
      "#endif\n",
      "__forceinline__ __device__ unsigned int\n",
      "cast_smem_ptr_to_int(const void* const smem_ptr)\n",
      "{\n",
      "  unsigned int smem_int;\n",
      "  asm volatile (\"{ .reg .u64 smem_int; cvta.to.shared.u64 smem_int, %1; cvt.u32.u64 %0, smem_int; }\"\n",
      "    : \"=r\"(smem_int) : \"l\"(smem_ptr));\n",
      "  return smem_int;\n",
      "}\n",
      "\n",
      "#if (((__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ >= 4)) || \\\n",
      "     (__CUDACC_VER_MAJOR__ > 11))\n",
      "#define TVM_ENABLE_L2_PREFETCH 1\n",
      "#else\n",
      "#define TVM_ENABLE_L2_PREFETCH 0\n",
      "#endif\n",
      "\n",
      "#ifdef _WIN32\n",
      "  using uint = unsigned int;\n",
      "  using uchar = unsigned char;\n",
      "  using ushort = unsigned short;\n",
      "  using int64_t = long long;\n",
      "  using uint64_t = unsigned long long;\n",
      "#else\n",
      "  #define uint unsigned int\n",
      "  #define uchar unsigned char\n",
      "  #define ushort unsigned short\n",
      "  #define int64_t long long\n",
      "  #define uint64_t unsigned long long\n",
      "#endif\n",
      "\n",
      "#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ == 800) \n",
      "#define TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST 1\n",
      "#else\n",
      "#define TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST 0\n",
      "#endif\n",
      "extern \"C\" __global__ void __launch_bounds__(32) matmul_n1024k1024_f16xi4_tcx16x128x64w16x32_opt_m_16(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m);\n",
      "extern \"C\" __global__ void __launch_bounds__(32) matmul_n1024k1024_f16xi4_tcx16x64x128w16x16_opt_m_32(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m);\n",
      "extern \"C\" __global__ void __launch_bounds__(128) matmul_n1024k1024_f16xi4_simt_opt_m_1(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m);\n",
      "extern \"C\" __global__ void __launch_bounds__(32) matmul_n1024k1024_f16xi4_tcx32x64x128w16x32_opt_m_128(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m);\n",
      "extern \"C\" __global__ void __launch_bounds__(128) matmul_n1024k1024_f16xi4_tcx32x64x128w16x32_opt_m_256(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m);\n",
      "extern \"C\" __global__ void __launch_bounds__(32) matmul_n1024k1024_f16xi4_tcx32x16x256w16x16_opt_m_64(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m);\n",
      "extern \"C\" __global__ void __launch_bounds__(32) matmul_n1024k1024_f16xi4_tcx16x128x64w16x32_opt_m_16(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m) {\n",
      "  extern __shared__ uchar buf_dyn_shmem[];\n",
      "  half C_reindex_pad_shared_dyn_warp[8];\n",
      "  signed char B_local[4];\n",
      "  half B_decode_reindex_local[8];\n",
      "  half A_reindex_pad_shared_dyn_warp[8];\n",
      "  half B_decode_reindex_shared_dyn_warp[8];\n",
      "  for (int var = 0; var < 1; ++var) {\n",
      "    for (int i = 0; i < 8; ++i) {\n",
      "C_reindex_pad_shared_dyn_warp[0 + i] = 0.0;}\n",
      ";\n",
      "    for (int ax3_0_0 = 0; ax3_0_0 < 4; ++ax3_0_0) {\n",
      "      __syncthreads();\n",
      "      #pragma unroll\n",
      "      for (int ax0_ax1_ax2_fused_0 = 0; ax0_ax1_ax2_fused_0 < 16; ++ax0_ax1_ax2_fused_0) {\n",
      "        uint4 condval;\n",
      "        if ((((((int)blockIdx.y) * 16) + ax0_ax1_ax2_fused_0) < m)) {\n",
      "          condval = *(uint4*)(A + ((((((int)blockIdx.y) * 16384) + (ax0_ax1_ax2_fused_0 * 1024)) + (ax3_0_0 * 256)) + (((int)threadIdx.x) * 8)));\n",
      "        } else {\n",
      "          condval = make_uint4(__pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)));\n",
      "        }\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((ax0_ax1_ax2_fused_0 * 264) + (((int)threadIdx.x) * 8)) + 256)) = condval;\n",
      "      }\n",
      "      for (int ax1_ax2_0_fused_0 = 0; ax1_ax2_0_fused_0 < 16; ++ax1_ax2_0_fused_0) {\n",
      "        *(int*)(B_local + 0) = *(int*)(B + ((((((int)blockIdx.x) * 8192) + (ax1_ax2_0_fused_0 * 512)) + (ax3_0_0 * 128)) + (((int)threadIdx.x) * 4)));\n",
      "        decode_i4s_to_f16((&(B_local[0])), (&(B_decode_reindex_local[0])), 8);\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((ax1_ax2_0_fused_0 * 264) + (((int)threadIdx.x) * 8)) + 4480)) = *(uint4*)(B_decode_reindex_local + 0);\n",
      "      }\n",
      "      __syncthreads();\n",
      "      for (int ax3_0_1 = 0; ax3_0_1 < 16; ++ax3_0_1) {\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 256)])) + (((((int)threadIdx.x) & 15) * 264) + ((((int)threadIdx.x) >> 4) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 256)])) + (((((int)threadIdx.x) & 15) * 264) + ((((int)threadIdx.x) >> 4) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 4480)])) + ((((((int)threadIdx.x) >> 4) * 2112) + ((((int)threadIdx.x) & 7) * 264)) + (((((int)threadIdx.x) & 15) >> 3) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 4480)])) + ((((((int)threadIdx.x) >> 4) * 2112) + ((((int)threadIdx.x) & 7) * 264)) + (((((int)threadIdx.x) & 15) >> 3) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1]));\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1]));\n",
      "  }\n",
      "      }\n",
      "    }\n",
      "    __syncthreads();\n",
      "    for (int local_id = 0; local_id < 8; local_id+=2) {\n",
      "*((uint *)&(&(((half*)buf_dyn_shmem)[0]))[((((((local_id % 4) / 2) * 8) + (threadIdx.x / 4)) * 16) + ((((local_id / 4) * 8) + ((threadIdx.x % 4) * 2)) + (local_id % 2)))]) = *((uint *)&C_reindex_pad_shared_dyn_warp[0 + local_id]);\n",
      "}\n",
      ";\n",
      "  }\n",
      "  __syncthreads();\n",
      "  #pragma unroll\n",
      "  for (int ax0_ax1_ax2_ax3_ax4_fused_0 = 0; ax0_ax1_ax2_ax3_ax4_fused_0 < 1; ++ax0_ax1_ax2_ax3_ax4_fused_0) {\n",
      "    if (((((int)blockIdx.y) * 16) + (((int)threadIdx.x) >> 1)) < m) {\n",
      "      *(uint4*)(C + ((((((int)blockIdx.y) * 16384) + ((((int)threadIdx.x) >> 1) * 1024)) + (((int)blockIdx.x) * 16)) + ((((int)threadIdx.x) & 1) * 8))) = *(uint4*)(((half*)buf_dyn_shmem) + (((int)threadIdx.x) * 8));\n",
      "    }\n",
      "  }\n",
      "}\n",
      "\n",
      "extern \"C\" __global__ void __launch_bounds__(32) matmul_n1024k1024_f16xi4_tcx16x64x128w16x16_opt_m_32(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m) {\n",
      "  extern __shared__ uchar buf_dyn_shmem[];\n",
      "  half C_reindex_pad_shared_dyn_warp[8];\n",
      "  signed char B_local[4];\n",
      "  half B_decode_reindex_local[8];\n",
      "  half A_reindex_pad_shared_dyn_warp[8];\n",
      "  half B_decode_reindex_shared_dyn_warp[8];\n",
      "  for (int var = 0; var < 1; ++var) {\n",
      "    for (int i = 0; i < 8; ++i) {\n",
      "C_reindex_pad_shared_dyn_warp[0 + i] = 0.0;}\n",
      ";\n",
      "    for (int ax3_0_0 = 0; ax3_0_0 < 4; ++ax3_0_0) {\n",
      "      __syncthreads();\n",
      "      #pragma unroll\n",
      "      for (int ax0_ax1_ax2_fused_0 = 0; ax0_ax1_ax2_fused_0 < 16; ++ax0_ax1_ax2_fused_0) {\n",
      "        uint4 condval;\n",
      "        if ((((((int)blockIdx.y) * 16) + ax0_ax1_ax2_fused_0) < m)) {\n",
      "          condval = *(uint4*)(A + ((((((int)blockIdx.y) * 16384) + (ax0_ax1_ax2_fused_0 * 1024)) + (ax3_0_0 * 256)) + (((int)threadIdx.x) * 8)));\n",
      "        } else {\n",
      "          condval = make_uint4(__pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)));\n",
      "        }\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((ax0_ax1_ax2_fused_0 * 264) + (((int)threadIdx.x) * 8)) + 256)) = condval;\n",
      "      }\n",
      "      for (int ax1_ax2_0_fused_0 = 0; ax1_ax2_0_fused_0 < 16; ++ax1_ax2_0_fused_0) {\n",
      "        *(int*)(B_local + 0) = *(int*)(B + ((((((int)blockIdx.x) * 8192) + (ax1_ax2_0_fused_0 * 512)) + (ax3_0_0 * 128)) + (((int)threadIdx.x) * 4)));\n",
      "        decode_i4s_to_f16((&(B_local[0])), (&(B_decode_reindex_local[0])), 8);\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((ax1_ax2_0_fused_0 * 264) + (((int)threadIdx.x) * 8)) + 4480)) = *(uint4*)(B_decode_reindex_local + 0);\n",
      "      }\n",
      "      __syncthreads();\n",
      "      for (int ax3_0_1 = 0; ax3_0_1 < 16; ++ax3_0_1) {\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 256)])) + (((((int)threadIdx.x) & 15) * 264) + ((((int)threadIdx.x) >> 4) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 256)])) + (((((int)threadIdx.x) & 15) * 264) + ((((int)threadIdx.x) >> 4) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 4480)])) + ((((((int)threadIdx.x) >> 4) * 2112) + ((((int)threadIdx.x) & 7) * 264)) + (((((int)threadIdx.x) & 15) >> 3) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 4480)])) + ((((((int)threadIdx.x) >> 4) * 2112) + ((((int)threadIdx.x) & 7) * 264)) + (((((int)threadIdx.x) & 15) >> 3) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1]));\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1]));\n",
      "  }\n",
      "      }\n",
      "    }\n",
      "    __syncthreads();\n",
      "    for (int local_id = 0; local_id < 8; local_id+=2) {\n",
      "*((uint *)&(&(((half*)buf_dyn_shmem)[0]))[((((((local_id % 4) / 2) * 8) + (threadIdx.x / 4)) * 16) + ((((local_id / 4) * 8) + ((threadIdx.x % 4) * 2)) + (local_id % 2)))]) = *((uint *)&C_reindex_pad_shared_dyn_warp[0 + local_id]);\n",
      "}\n",
      ";\n",
      "  }\n",
      "  __syncthreads();\n",
      "  #pragma unroll\n",
      "  for (int ax0_ax1_ax2_ax3_ax4_fused_0 = 0; ax0_ax1_ax2_ax3_ax4_fused_0 < 1; ++ax0_ax1_ax2_ax3_ax4_fused_0) {\n",
      "    if (((((int)blockIdx.y) * 16) + (((int)threadIdx.x) >> 1)) < m) {\n",
      "      *(uint4*)(C + ((((((int)blockIdx.y) * 16384) + ((((int)threadIdx.x) >> 1) * 1024)) + (((int)blockIdx.x) * 16)) + ((((int)threadIdx.x) & 1) * 8))) = *(uint4*)(((half*)buf_dyn_shmem) + (((int)threadIdx.x) * 8));\n",
      "    }\n",
      "  }\n",
      "}\n",
      "\n",
      "extern \"C\" __global__ void __launch_bounds__(128) matmul_n1024k1024_f16xi4_simt_opt_m_1(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m) {\n",
      "  half in_thread_C_local[1];\n",
      "  signed char B_local[4];\n",
      "  half B_decode_local[8];\n",
      "  half A_local[8];\n",
      "  __shared__ half red_result[2];\n",
      "  in_thread_C_local[0] = __float2half_rn(0.000000e+00f);\n",
      "  for (int ax2_0 = 0; ax2_0 < 2; ++ax2_0) {\n",
      "    *(int*)(B_local + 0) = *(int*)(B + ((((((int)blockIdx.x) * 1024) + (((int)threadIdx.y) * 512)) + (ax2_0 * 256)) + (((int)threadIdx.x) * 4)));\n",
      "    decode_i4s_to_f16((&(B_local[0])), (&(B_decode_local[0])), 8);\n",
      "    *(uint4*)(A_local + 0) = *(uint4*)(A + (((((int)blockIdx.z) * 1024) + (ax2_0 * 512)) + (((int)threadIdx.x) * 8)));\n",
      "    for (int ax2_2_0 = 0; ax2_2_0 < 4; ++ax2_2_0) {\n",
      "      for (int ax2_2_1 = 0; ax2_2_1 < 2; ++ax2_2_1) {\n",
      "        in_thread_C_local[0] = (in_thread_C_local[0] + (A_local[((ax2_2_0 * 2) + ax2_2_1)] * B_decode_local[((ax2_2_0 * 2) + ax2_2_1)]));\n",
      "      }\n",
      "    }\n",
      "  }\n",
      "  half red_buf0[1];\n",
      "  uint mask[1];\n",
      "  half t0[1];\n",
      "  half red_buf0_1[1];\n",
      "  uint mask_1[1];\n",
      "  half t0_1[1];\n",
      "  __shared__ half red_buf_staging[4];\n",
      "  red_buf0_1[0] = in_thread_C_local[0];\n",
      "  mask_1[0] = __activemask();\n",
      "  t0_1[0] = __shfl_down_sync(mask_1[0], red_buf0_1[0], 16, 32);\n",
      "  red_buf0_1[0] = (red_buf0_1[0] + t0_1[0]);\n",
      "  t0_1[0] = __shfl_down_sync(mask_1[0], red_buf0_1[0], 8, 32);\n",
      "  red_buf0_1[0] = (red_buf0_1[0] + t0_1[0]);\n",
      "  t0_1[0] = __shfl_down_sync(mask_1[0], red_buf0_1[0], 4, 32);\n",
      "  red_buf0_1[0] = (red_buf0_1[0] + t0_1[0]);\n",
      "  t0_1[0] = __shfl_down_sync(mask_1[0], red_buf0_1[0], 2, 32);\n",
      "  red_buf0_1[0] = (red_buf0_1[0] + t0_1[0]);\n",
      "  t0_1[0] = __shfl_down_sync(mask_1[0], red_buf0_1[0], 1, 32);\n",
      "  red_buf0_1[0] = (red_buf0_1[0] + t0_1[0]);\n",
      "  if ((((int)threadIdx.x) % 32) == 0) {\n",
      "    red_buf_staging[((((int)threadIdx.y) * 2) + (((int)threadIdx.x) >> 5))] = red_buf0_1[0];\n",
      "  }\n",
      "  __syncthreads();\n",
      "  if (((int)threadIdx.x) < 2) {\n",
      "    red_buf0[0] = red_buf_staging[((((int)threadIdx.y) * 2) + ((int)threadIdx.x))];\n",
      "  }\n",
      "  mask[0] = __activemask();\n",
      "  t0[0] = __shfl_down_sync(mask[0], red_buf0[0], 1, 32);\n",
      "  red_buf0[0] = (red_buf0[0] + t0[0]);\n",
      "  if (((int)threadIdx.x) == 0) {\n",
      "    ((volatile half*)red_result)[((int)threadIdx.y)] = red_buf0[0];\n",
      "  }\n",
      "  __syncthreads();\n",
      "  if (((int)threadIdx.x) == 0) {\n",
      "    C[(((((int)blockIdx.z) * 1024) + (((int)blockIdx.x) * 2)) + ((int)threadIdx.y))] = (half)(((volatile half*)red_result)[((int)threadIdx.y)]);\n",
      "  }\n",
      "}\n",
      "\n",
      "extern \"C\" __global__ void __launch_bounds__(32) matmul_n1024k1024_f16xi4_tcx32x64x128w16x32_opt_m_128(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m) {\n",
      "  extern __shared__ uchar buf_dyn_shmem[];\n",
      "  half C_reindex_pad_shared_dyn_warp[8];\n",
      "  signed char B_local[4];\n",
      "  half B_decode_reindex_local[8];\n",
      "  half A_reindex_pad_shared_dyn_warp[8];\n",
      "  half B_decode_reindex_shared_dyn_warp[8];\n",
      "  for (int var = 0; var < 1; ++var) {\n",
      "    for (int i = 0; i < 8; ++i) {\n",
      "C_reindex_pad_shared_dyn_warp[0 + i] = 0.0;}\n",
      ";\n",
      "    for (int ax3_0_0 = 0; ax3_0_0 < 4; ++ax3_0_0) {\n",
      "      __syncthreads();\n",
      "      #pragma unroll\n",
      "      for (int ax0_ax1_ax2_fused_0 = 0; ax0_ax1_ax2_fused_0 < 16; ++ax0_ax1_ax2_fused_0) {\n",
      "        uint4 condval;\n",
      "        if ((((((int)blockIdx.y) * 16) + ax0_ax1_ax2_fused_0) < m)) {\n",
      "          condval = *(uint4*)(A + ((((((int)blockIdx.y) * 16384) + (ax0_ax1_ax2_fused_0 * 1024)) + (ax3_0_0 * 256)) + (((int)threadIdx.x) * 8)));\n",
      "        } else {\n",
      "          condval = make_uint4(__pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)));\n",
      "        }\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((ax0_ax1_ax2_fused_0 * 264) + (((int)threadIdx.x) * 8)) + 256)) = condval;\n",
      "      }\n",
      "      for (int ax1_ax2_0_fused_0 = 0; ax1_ax2_0_fused_0 < 16; ++ax1_ax2_0_fused_0) {\n",
      "        *(int*)(B_local + 0) = *(int*)(B + ((((((int)blockIdx.x) * 8192) + (ax1_ax2_0_fused_0 * 512)) + (ax3_0_0 * 128)) + (((int)threadIdx.x) * 4)));\n",
      "        decode_i4s_to_f16((&(B_local[0])), (&(B_decode_reindex_local[0])), 8);\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((ax1_ax2_0_fused_0 * 264) + (((int)threadIdx.x) * 8)) + 4480)) = *(uint4*)(B_decode_reindex_local + 0);\n",
      "      }\n",
      "      __syncthreads();\n",
      "      for (int ax3_0_1 = 0; ax3_0_1 < 16; ++ax3_0_1) {\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 256)])) + (((((int)threadIdx.x) & 15) * 264) + ((((int)threadIdx.x) >> 4) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 256)])) + (((((int)threadIdx.x) & 15) * 264) + ((((int)threadIdx.x) >> 4) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 4480)])) + ((((((int)threadIdx.x) >> 4) * 2112) + ((((int)threadIdx.x) & 7) * 264)) + (((((int)threadIdx.x) & 15) >> 3) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 4480)])) + ((((((int)threadIdx.x) >> 4) * 2112) + ((((int)threadIdx.x) & 7) * 264)) + (((((int)threadIdx.x) & 15) >> 3) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1]));\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1]));\n",
      "  }\n",
      "      }\n",
      "    }\n",
      "    __syncthreads();\n",
      "    for (int local_id = 0; local_id < 8; local_id+=2) {\n",
      "*((uint *)&(&(((half*)buf_dyn_shmem)[0]))[((((((local_id % 4) / 2) * 8) + (threadIdx.x / 4)) * 16) + ((((local_id / 4) * 8) + ((threadIdx.x % 4) * 2)) + (local_id % 2)))]) = *((uint *)&C_reindex_pad_shared_dyn_warp[0 + local_id]);\n",
      "}\n",
      ";\n",
      "  }\n",
      "  __syncthreads();\n",
      "  #pragma unroll\n",
      "  for (int ax0_ax1_ax2_ax3_ax4_fused_0 = 0; ax0_ax1_ax2_ax3_ax4_fused_0 < 1; ++ax0_ax1_ax2_ax3_ax4_fused_0) {\n",
      "    if (((((int)blockIdx.y) * 16) + (((int)threadIdx.x) >> 1)) < m) {\n",
      "      *(uint4*)(C + ((((((int)blockIdx.y) * 16384) + ((((int)threadIdx.x) >> 1) * 1024)) + (((int)blockIdx.x) * 16)) + ((((int)threadIdx.x) & 1) * 8))) = *(uint4*)(((half*)buf_dyn_shmem) + (((int)threadIdx.x) * 8));\n",
      "    }\n",
      "  }\n",
      "}\n",
      "\n",
      "extern \"C\" __global__ void __launch_bounds__(128) matmul_n1024k1024_f16xi4_tcx32x64x128w16x32_opt_m_256(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m) {\n",
      "  extern __shared__ uchar buf_dyn_shmem[];\n",
      "  half C_reindex_pad_shared_dyn_warp[8];\n",
      "  signed char B_local[4];\n",
      "  half B_decode_reindex_local[8];\n",
      "  half A_reindex_pad_shared_dyn_warp[8];\n",
      "  half B_decode_reindex_shared_dyn_warp[8];\n",
      "  for (int var = 0; var < 1; ++var) {\n",
      "    for (int i = 0; i < 8; ++i) {\n",
      "C_reindex_pad_shared_dyn_warp[0 + i] = 0.0;}\n",
      ";\n",
      "    for (int ax3_0_0 = 0; ax3_0_0 < 8; ++ax3_0_0) {\n",
      "      __syncthreads();\n",
      "      #pragma unroll\n",
      "      for (int ax0_ax1_ax2_fused_0 = 0; ax0_ax1_ax2_fused_0 < 2; ++ax0_ax1_ax2_fused_0) {\n",
      "        uint4 condval;\n",
      "        if ((((((((int)blockIdx.y) * 16) + (ax0_ax1_ax2_fused_0 * 8)) + (((int)threadIdx.z) * 2)) + (((int)threadIdx.x) >> 4)) < m)) {\n",
      "          condval = *(uint4*)(A + ((((((((int)blockIdx.y) * 16384) + (ax0_ax1_ax2_fused_0 * 8192)) + (((int)threadIdx.z) * 2048)) + ((((int)threadIdx.x) >> 4) * 1024)) + (ax3_0_0 * 128)) + ((((int)threadIdx.x) & 15) * 8)));\n",
      "        } else {\n",
      "          condval = make_uint4(__pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)));\n",
      "        }\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + ((((ax0_ax1_ax2_fused_0 * 1088) + (((int)threadIdx.z) * 272)) + ((((int)threadIdx.x) >> 4) * 136)) + ((((int)threadIdx.x) & 15) * 8))) = condval;\n",
      "      }\n",
      "      for (int ax1_ax2_0_fused_0 = 0; ax1_ax2_0_fused_0 < 8; ++ax1_ax2_0_fused_0) {\n",
      "        *(int*)(B_local + 0) = *(int*)(B + ((((((((int)blockIdx.x) * 32768) + (ax1_ax2_0_fused_0 * 4096)) + (((int)threadIdx.z) * 1024)) + ((((int)threadIdx.x) >> 4) * 512)) + (ax3_0_0 * 64)) + ((((int)threadIdx.x) & 15) * 4)));\n",
      "        decode_i4s_to_f16((&(B_local[0])), (&(B_decode_reindex_local[0])), 8);\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((((ax1_ax2_0_fused_0 * 1088) + (((int)threadIdx.z) * 272)) + ((((int)threadIdx.x) >> 4) * 136)) + ((((int)threadIdx.x) & 15) * 8)) + 2176)) = *(uint4*)(B_decode_reindex_local + 0);\n",
      "      }\n",
      "      __syncthreads();\n",
      "      for (int ax3_0_1 = 0; ax3_0_1 < 8; ++ax3_0_1) {\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[(ax3_0_1 * 16)])) + (((((int)threadIdx.x) & 15) * 136) + ((((int)threadIdx.x) >> 4) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[(ax3_0_1 * 16)])) + (((((int)threadIdx.x) & 15) * 136) + ((((int)threadIdx.x) >> 4) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[(((((int)threadIdx.z) * 2176) + (ax3_0_1 * 16)) + 2176)])) + ((((((int)threadIdx.x) >> 4) * 1088) + ((((int)threadIdx.x) & 7) * 136)) + (((((int)threadIdx.x) & 15) >> 3) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[(((((int)threadIdx.z) * 2176) + (ax3_0_1 * 16)) + 2176)])) + ((((((int)threadIdx.x) >> 4) * 1088) + ((((int)threadIdx.x) & 7) * 136)) + (((((int)threadIdx.x) & 15) >> 3) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1]));\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1]));\n",
      "  }\n",
      "      }\n",
      "    }\n",
      "    __syncthreads();\n",
      "    for (int local_id = 0; local_id < 8; local_id+=2) {\n",
      "*((uint *)&(&(((half*)buf_dyn_shmem)[(((int)threadIdx.z) * 256)]))[((((((local_id % 4) / 2) * 8) + (threadIdx.x / 4)) * 16) + ((((local_id / 4) * 8) + ((threadIdx.x % 4) * 2)) + (local_id % 2)))]) = *((uint *)&C_reindex_pad_shared_dyn_warp[0 + local_id]);\n",
      "}\n",
      ";\n",
      "  }\n",
      "  __syncthreads();\n",
      "  #pragma unroll\n",
      "  for (int ax0_ax1_ax2_ax3_ax4_fused_0 = 0; ax0_ax1_ax2_ax3_ax4_fused_0 < 4; ++ax0_ax1_ax2_ax3_ax4_fused_0) {\n",
      "    if (((((int)blockIdx.y) * 16) + (((int)threadIdx.x) >> 1)) < m) {\n",
      "      *(uint4*)(C + (((((((int)blockIdx.y) * 16384) + ((((int)threadIdx.x) >> 1) * 1024)) + (((int)blockIdx.x) * 64)) + (ax0_ax1_ax2_ax3_ax4_fused_0 * 16)) + ((((int)threadIdx.x) & 1) * 8))) = *(uint4*)(((half*)buf_dyn_shmem) + ((ax0_ax1_ax2_ax3_ax4_fused_0 * 256) + (((int)threadIdx.x) * 8)));\n",
      "    }\n",
      "  }\n",
      "}\n",
      "\n",
      "extern \"C\" __global__ void __launch_bounds__(32) matmul_n1024k1024_f16xi4_tcx32x16x256w16x16_opt_m_64(half* __restrict__ A, signed char* __restrict__ B, half* __restrict__ C, int m) {\n",
      "  extern __shared__ uchar buf_dyn_shmem[];\n",
      "  half C_reindex_pad_shared_dyn_warp[8];\n",
      "  signed char B_local[4];\n",
      "  half B_decode_reindex_local[8];\n",
      "  half A_reindex_pad_shared_dyn_warp[8];\n",
      "  half B_decode_reindex_shared_dyn_warp[8];\n",
      "  for (int var = 0; var < 1; ++var) {\n",
      "    for (int i = 0; i < 8; ++i) {\n",
      "C_reindex_pad_shared_dyn_warp[0 + i] = 0.0;}\n",
      ";\n",
      "    for (int ax3_0_0 = 0; ax3_0_0 < 4; ++ax3_0_0) {\n",
      "      __syncthreads();\n",
      "      #pragma unroll\n",
      "      for (int ax0_ax1_ax2_fused_0 = 0; ax0_ax1_ax2_fused_0 < 16; ++ax0_ax1_ax2_fused_0) {\n",
      "        uint4 condval;\n",
      "        if ((((((int)blockIdx.y) * 16) + ax0_ax1_ax2_fused_0) < m)) {\n",
      "          condval = *(uint4*)(A + ((((((int)blockIdx.y) * 16384) + (ax0_ax1_ax2_fused_0 * 1024)) + (ax3_0_0 * 256)) + (((int)threadIdx.x) * 8)));\n",
      "        } else {\n",
      "          condval = make_uint4(__pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)), __pack_half2(__float2half_rn(0.000000e+00f), __float2half_rn(0.000000e+00f)));\n",
      "        }\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((ax0_ax1_ax2_fused_0 * 264) + (((int)threadIdx.x) * 8)) + 256)) = condval;\n",
      "      }\n",
      "      for (int ax1_ax2_0_fused_0 = 0; ax1_ax2_0_fused_0 < 16; ++ax1_ax2_0_fused_0) {\n",
      "        *(int*)(B_local + 0) = *(int*)(B + ((((((int)blockIdx.x) * 8192) + (ax1_ax2_0_fused_0 * 512)) + (ax3_0_0 * 128)) + (((int)threadIdx.x) * 4)));\n",
      "        decode_i4s_to_f16((&(B_local[0])), (&(B_decode_reindex_local[0])), 8);\n",
      "        *(uint4*)(((half*)buf_dyn_shmem) + (((ax1_ax2_0_fused_0 * 264) + (((int)threadIdx.x) * 8)) + 4480)) = *(uint4*)(B_decode_reindex_local + 0);\n",
      "      }\n",
      "      __syncthreads();\n",
      "      for (int ax3_0_1 = 0; ax3_0_1 < 16; ++ax3_0_1) {\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 256)])) + (((((int)threadIdx.x) & 15) * 264) + ((((int)threadIdx.x) >> 4) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 256)])) + (((((int)threadIdx.x) & 15) * 264) + ((((int)threadIdx.x) >> 4) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(A_reindex_pad_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    unsigned int addr;\n",
      "#if TVM_ENBALE_EFFICIENT_SMEM_PTR_CAST\n",
      "    addr = static_cast<unsigned int>(__cvta_generic_to_shared((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 4480)])) + ((((((int)threadIdx.x) >> 4) * 2112) + ((((int)threadIdx.x) & 7) * 264)) + (((((int)threadIdx.x) & 15) >> 3) * 8)))));\n",
      "#else\n",
      "    __asm__ __volatile__(\n",
      "      \"{ .reg .u64 addr; cvta.to.shared.u64 addr, %1; cvt.u32.u64 %0, addr; }\\n\"\n",
      "      : \"=r\"(addr)\n",
      "      : \"l\"((void *)((&(((half*)buf_dyn_shmem)[((ax3_0_1 * 16) + 4480)])) + ((((((int)threadIdx.x) >> 4) * 2112) + ((((int)threadIdx.x) & 7) * 264)) + (((((int)threadIdx.x) & 15) >> 3) * 8))))\n",
      "    );\n",
      "#endif\n",
      "    __asm__ __volatile__(\n",
      "      \"ldmatrix.sync.aligned.m8n8.x4.shared.b16\"\n",
      "      \"{%0, %1, %2, %3}, [%4];\\n\"\n",
      "      : \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[1]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[2]), \"=r\"(((unsigned *)(B_decode_reindex_shared_dyn_warp + 0))[3])\n",
      "      : \"r\"(addr)\n",
      "    );\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 0))[1]));\n",
      "  }\n",
      "\n",
      "  {\n",
      "    __asm__ __volatile__(\n",
      "      \"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16\"\n",
      "      \"{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\\n\"\n",
      "      :  \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"=r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1])\n",
      "      : \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[0]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[1]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[2]), \"r\"(((unsigned *)((half*)A_reindex_pad_shared_dyn_warp + 0))[3]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)((half*)B_decode_reindex_shared_dyn_warp + 4))[1]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[0]), \"r\"(((unsigned *)(C_reindex_pad_shared_dyn_warp + 4))[1]));\n",
      "  }\n",
      "      }\n",
      "    }\n",
      "    __syncthreads();\n",
      "    for (int local_id = 0; local_id < 8; local_id+=2) {\n",
      "*((uint *)&(&(((half*)buf_dyn_shmem)[0]))[((((((local_id % 4) / 2) * 8) + (threadIdx.x / 4)) * 16) + ((((local_id / 4) * 8) + ((threadIdx.x % 4) * 2)) + (local_id % 2)))]) = *((uint *)&C_reindex_pad_shared_dyn_warp[0 + local_id]);\n",
      "}\n",
      ";\n",
      "  }\n",
      "  __syncthreads();\n",
      "  #pragma unroll\n",
      "  for (int ax0_ax1_ax2_ax3_ax4_fused_0 = 0; ax0_ax1_ax2_ax3_ax4_fused_0 < 1; ++ax0_ax1_ax2_ax3_ax4_fused_0) {\n",
      "    if (((((int)blockIdx.y) * 16) + (((int)threadIdx.x) >> 1)) < m) {\n",
      "      *(uint4*)(C + ((((((int)blockIdx.y) * 16384) + ((((int)threadIdx.x) >> 1) * 1024)) + (((int)blockIdx.x) * 16)) + ((((int)threadIdx.x) & 1) * 8))) = *(uint4*)(((half*)buf_dyn_shmem) + (((int)threadIdx.x) * 8));\n",
      "    }\n",
      "  }\n",
      "}\n",
      "\n",
      "\n",
      "extern \"C\" void init() {\n",
      "    \n",
      "    cudaFuncSetAttribute(matmul_n1024k1024_f16xi4_tcx16x128x64w16x32_opt_m_16, cudaFuncAttributeMaxDynamicSharedMemorySize, 17408);\n",
      "\n",
      "    cudaFuncSetAttribute(matmul_n1024k1024_f16xi4_tcx16x64x128w16x16_opt_m_32, cudaFuncAttributeMaxDynamicSharedMemorySize, 17408);\n",
      "\n",
      "    cudaFuncSetAttribute(matmul_n1024k1024_f16xi4_tcx32x64x128w16x32_opt_m_128, cudaFuncAttributeMaxDynamicSharedMemorySize, 17408);\n",
      "\n",
      "    cudaFuncSetAttribute(matmul_n1024k1024_f16xi4_tcx32x64x128w16x32_opt_m_256, cudaFuncAttributeMaxDynamicSharedMemorySize, 21760);\n",
      "\n",
      "    cudaFuncSetAttribute(matmul_n1024k1024_f16xi4_tcx32x16x256w16x16_opt_m_64, cudaFuncAttributeMaxDynamicSharedMemorySize, 17408);\n",
      "\n",
      "}\n",
      "\n",
      "extern \"C\" void call(half* __restrict__ A, int8_t* __restrict__ B, half* __restrict__ C, int m, cudaStream_t stream=cudaStreamDefault) {\n",
      "  if (m == 0) return; \n",
      "  if (m <= 1) {\n",
      "    matmul_n1024k1024_f16xi4_simt_opt_m_1<<<dim3(512, 1, m), dim3(64, 2, 1), 0, stream>>>(A, B, C, m); \n",
      "  }\n",
      "  else if (m <= 16) {\n",
      "    matmul_n1024k1024_f16xi4_tcx16x128x64w16x32_opt_m_16<<<dim3(64, (m + 15) / 16, 1), dim3(32, 1, 1), 17408, stream>>>(A, B, C, m); \n",
      "  }\n",
      "  else if (m <= 32) {\n",
      "    matmul_n1024k1024_f16xi4_tcx16x64x128w16x16_opt_m_32<<<dim3(64, (m + 15) / 16, 1), dim3(32, 1, 1), 17408, stream>>>(A, B, C, m); \n",
      "  }\n",
      "  else if (m <= 64) {\n",
      "    matmul_n1024k1024_f16xi4_tcx32x16x256w16x16_opt_m_64<<<dim3(64, (m + 15) / 16, 1), dim3(32, 1, 1), 17408, stream>>>(A, B, C, m); \n",
      "  }\n",
      "  else if (m <= 128) {\n",
      "    matmul_n1024k1024_f16xi4_tcx32x64x128w16x32_opt_m_128<<<dim3(64, (m + 15) / 16, 1), dim3(32, 1, 1), 17408, stream>>>(A, B, C, m); \n",
      "  }\n",
      "  else if (m <= 256) {\n",
      "    matmul_n1024k1024_f16xi4_tcx32x64x128w16x32_opt_m_256<<<dim3(16, (m + 15) / 16, 1), dim3(32, 1, 4), 21760, stream>>>(A, B, C, m); \n",
      "  }\n",
      "  else {\n",
      "    matmul_n1024k1024_f16xi4_tcx32x64x128w16x32_opt_m_256<<<dim3(16, (m + 15) / 16, 1), dim3(32, 1, 4), 21760, stream>>>(A, B, C, m); \n",
      "  }\n",
      "\n",
      "}\n",
      "\n"
     ]
    }
   ],
   "source": [
    "matmul.hardware_aware_finetune(topk=20, parallel_build=True)\n",
    "print(matmul.get_source())"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3 (ipykernel)",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.10.14"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 5
}
